/*
 *  cudaHost.h
 *  HPC
 *
 *  Created by Arun Purushan and Prabhat Aravind on 11/23/09.
 *  Copyright 2009 UFL. All rights reserved.
 *
 */

#define NUM_THREAD 256 //for a dataset of size 1500*1000
#define NUM_BLOCK (int)ceil(n/NUM_THREAD)
#include <cutil_inline.h>
#include "cuda_kernel.cu"
#include <stdio.h>
#include <time.h>
#include <sys/time.h>
//using namespace std
extern  void updateBC(int n, int k, int d, int n_start, int n_end);
void cudaComputeWPDF(double* answer1, double* answer2, double* answer3, double _newW1, double _newW2, double _newW3, int row, int col
				,double** _bw, int** _bc, int** _bg, double** _wMatrix, int n, int k, int d)
{

	int memSize_bc = 0, memSize_bg = 0, memSize_wMatrix = 0, memSize_w = 0, memSize_tp = 0;
	//  Allocating Memory for threadProduct

	double* threadProduct1 = (double*) malloc((n+1) * sizeof(double));
	double* threadProduct2 = (double*) malloc((n+1) * sizeof(double));
	double* threadProduct3 = (double*) malloc((n+1) * sizeof(double));
	

	// create a copy of w(row,:) and update it using proposed weight
	double updateRatio1 = (1 - _newW1)/(1 - _bw[row][col]);
	double updateRatio2 = (1 - _newW2)/(1 - _bw[row][col]);
	double updateRatio3 = (1 - _newW3)/(1 - _bw[row][col]);
	double* _w1 = new double[d]; double* _w2 = new double[d]; double* _w3 = new double[d];
	for (int j=0; j<d; j++){
		_w1[j] = updateRatio1 * _bw[row][j];
		_w2[j] = updateRatio2 * _bw[row][j];
		_w3[j] = updateRatio3 * _bw[row][j];
	}
	_w1[col] = _newW1; _w2[col] = _newW2; _w3[col] = _newW3;
	// allocate device(GPU) memory
	memSize_bc      = n * k * sizeof(int);
	memSize_bg      = n * d * sizeof(int);
//	memSize_wMatrix = k * d * sizeof(double);
	memSize_wMatrix = n * d * sizeof(double);
	memSize_w       = d * sizeof(double);
	memSize_tp      = (n+1) * sizeof(double);
	
	int* c_bg = 0;
        cutilSafeCall(cudaMalloc((void**) &c_bg, memSize_bg));

	int* c_bc = 0;
        cutilSafeCall(cudaMalloc((void**) &c_bc, memSize_bc));
				  
	double* c_wMatrix = 0;
	cutilSafeCall(cudaMalloc((void**) &c_wMatrix, memSize_wMatrix));
	
	double* c_w1 = 0;
	cutilSafeCall(cudaMalloc((void**) &c_w1, memSize_w));
	double* c_w2 = 0;
	cutilSafeCall(cudaMalloc((void**) &c_w2, memSize_w));
	double* c_w3 = 0;
	cutilSafeCall(cudaMalloc((void**) &c_w3, memSize_w));
	
	
	double* c_tp1 = 0;
	cutilSafeCall(cudaMalloc((void**) &c_tp1, memSize_tp));
	double* c_tp2 = 0;
	cutilSafeCall(cudaMalloc((void**) &c_tp2, memSize_tp));
	double* c_tp3 = 0;
	cutilSafeCall(cudaMalloc((void**) &c_tp3, memSize_tp));

	//Copying the arrays to continuous locations for gpu transfer- 
	int *lbc          = (int *) malloc(n*k*sizeof(int));
	int *lbg          = (int *) malloc(n*d*sizeof(int));
	double *lwMatrix = (double *) malloc(n*d*sizeof(double));

	for(int i=0; i<n; i++){
		for(int j = 0; j<k; j++)
			*(lbc+(i*k+j)) = _bc[i][j];
	}
	for(int i=0; i<n; i++){
		for(int j = 0; j<d; j++)
			*(lbg+(i*d+j)) = _bg[i][j];
	}
	for(int i=0; i<n; i++){
		for(int j = 0; j<d; j++)
			*(lwMatrix+(i*d+j)) = _wMatrix[i][j]; 
	}
/*
for (int i=0;i<n;i++){
for(int j=0;j<d;j++)
printf("%lf ",_wMatrix[i][j]);
printf("\n");} exit(1);*/
	
	//Copying parameters to the device
struct timeval c_start; gettimeofday(&c_start, NULL);
	cutilSafeCall(cudaMemcpy(c_bg, lbg, memSize_bg, cudaMemcpyHostToDevice) );

	cutilSafeCall(cudaMemcpy(c_bc, lbc, memSize_bc, cudaMemcpyHostToDevice) );
	cutilSafeCall(cudaMemcpy(c_wMatrix, lwMatrix, memSize_wMatrix, cudaMemcpyHostToDevice) );

	cutilSafeCall(cudaMemcpy(c_w1, _w1, memSize_w,cudaMemcpyHostToDevice) );
	cutilSafeCall(cudaMemcpy(c_w2, _w2, memSize_w,cudaMemcpyHostToDevice) );
	cutilSafeCall(cudaMemcpy(c_w3, _w3, memSize_w,cudaMemcpyHostToDevice) );

//cutilSafeCall(cudaMemcpy(c_tp1, threadProduct1, memSize_tp, cudaMemcpyHostToDevice) );
//cutilSafeCall(cudaMemcpy(c_tp2, threadProduct2, memSize_tp, cudaMemcpyHostToDevice) );
//cutilSafeCall(cudaMemcpy(c_tp3, threadProduct3, memSize_tp, cudaMemcpyHostToDevice) );
	//Set up Blocks and threads
	dim3 dimBlock(NUM_BLOCK);
	dim3 dimThread(NUM_THREAD);
	
	//Possibility of preprocessing bc to get an array with reference to wmatrix rows whose bc values are 1. Performance improvement tru ensuring that all threads 
	//have work to perform.
	//Excecute the Kernel

	cudaComputeKernel<<< dimBlock, dimThread >>>(c_tp1, c_tp2, c_tp3, c_w1, c_w2, c_w3, row, c_bc, c_bg, c_wMatrix, n, k, d, NUM_THREAD);
	cutilCheckMsg("Kernel execution failed");
	
	
<<<<<<< .mine
	cutilSafeCall(cudaMemcpy(threadProduct1, c_tp1, memSize_tp,cudaMemcpyDeviceToHost) );
	cutilSafeCall(cudaMemcpy(threadProduct2, c_tp2, memSize_tp,cudaMemcpyDeviceToHost) );
	cutilSafeCall(cudaMemcpy(threadProduct3, c_tp3, memSize_tp,cudaMemcpyDeviceToHost) );
//for (int i=0; i<n ; i++)
//printf("%f %f %f\n", threadProduct1[i],threadProduct2[i],threadProduct3[i]);
//exit(1);


=======
	cutilSafeCall(cudaMemcpy(threadProduct1, c_tp1, memSize_tp,cudaMemcpyDeviceToHost) );
	cutilSafeCall(cudaMemcpy(threadProduct2, c_tp2, memSize_tp,cudaMemcpyDeviceToHost) );
	cutilSafeCall(cudaMemcpy(threadProduct3, c_tp3, memSize_tp,cudaMemcpyDeviceToHost) );
for (int i=0; i<n ; i++)
printf("%f %f %f\n", threadProduct1[i],threadProduct2[i],threadProduct3[i]);
//exit(1);


>>>>>>> .r10
	//Answer1 = product of threadProduct
*answer1=0; *answer2=0; *answer3=0;
	for (int i=0; i<n ; i++) {
		*answer1 += threadProduct1[i];
		*answer2 += threadProduct2[i];
		*answer3 += threadProduct3[i];
	}
	//To free the allocated device memory
<<<<<<< .mine
printf("\n\n%f %f %f\n ", *answer1,*answer2, *answer3);
delete lbc;
delete lbg;
delete lwMatrix;
=======
printf("\n\n%f %f %f\n ", *answer1,*answer2, *answer3);
>>>>>>> .r10
delete threadProduct1;
delete threadProduct2;
delete threadProduct3;
	cudaFree(c_bg);
	cudaFree(c_bc);
	cudaFree(c_wMatrix);
	cudaFree(c_w1);
	cudaFree(c_w2);
	cudaFree(c_w3);
	cudaFree(c_tp1);
	cudaFree(c_tp2);
	cudaFree(c_tp3);
struct timeval c_end; gettimeofday(&c_end, NULL);
printf("Cuda Elapsed Time %lf",(double)(c_end.tv_sec) + c_end.tv_usec/1000000.0 - (double)(c_start.tv_sec) - c_start.tv_usec/1000000.0);
//exit(1);	
	//update active component indicators
//	updateBC( n, k, d, 0, n);
//exit(1);

}
